/**************************************************************************/
/* N-Reinas Paralelo DEd                                                  */
/* Parametros:								  */
/*	1- Tamano del tablero						  */
/*	2- Cantidad de filas para formar las combinaciones		  */
/**************************************************************************/
/**************************************************************************/

#include <stdio.h>
#define CudaSafeCall( err ) __cudaSafeCall( err, __FILE__, __LINE__ )
#define CudaCheckError()    __cudaCheckError( __FILE__, __LINE__ )

inline void __cudaSafeCall(cudaError err, const char *file, const int line) {
	if (cudaSuccess != err) {
		fprintf(stderr, "cudaSafeCall() failed at %s:%i : %s\n", file, line,
				cudaGetErrorString(err));
		exit(-1);
	}

	return;
}

inline void __cudaCheckError(const char *file, const int line) {
	cudaError err = cudaGetLastError();
	if (cudaSuccess != err) {
		fprintf(stderr, "cudaCheckError() failed at %s:%i : %s\n", file, line,
				cudaGetErrorString(err));
		exit(-1);
	}

	// More careful checking. However, this will affect performance.
	// Comment away if needed.
	err = cudaDeviceSynchronize();
	if (cudaSuccess != err) {
		fprintf(stderr, "cudaCheckError() with sync failed at %s:%i : %s\n",
				file, line, cudaGetErrorString(err));
		exit(-1);
	}

	return;
}

/****************************************************************/
/* Determina la cantidad de soluciones  por rotacion y simetria */
/****************************************************************/
//Varibles compartidas por los threads de cada bloque
__shared__ unsigned long C2[1024];
__shared__ unsigned long C4[1024];
__shared__ unsigned long C8[1024];

__device__ void Check(unsigned long *BOARD, unsigned long *BOARDE,
		unsigned long TOPBIT, unsigned long ENDBIT, unsigned long *BOARD2,
		unsigned long *BOARD1) {
	unsigned long *own, *you, bit, ptn;

	//90-degree rotation

	if (*BOARD2 == 1) {
		for (ptn = 2, own = BOARD + 1; own <= BOARDE; own++, ptn <<= 1) {
			bit = 1;
			for (you = BOARDE; *you != ptn && *own >= bit; you--)
				bit <<= 1;
			if (*own > bit)
				return;
			if (*own < bit)
				break;
		}
		if (own > BOARDE) {
			//COUNT2++;

			C2[threadIdx.x] = C2[threadIdx.x] + 2;

			//printf("2 \n");
			return;
		}
	}
	// 180-degree rotation
	if (*BOARDE == ENDBIT) {
		for (you = BOARDE - 1, own = BOARD + 1; own <= BOARDE; own++, you--) {
			bit = 1;
			for (ptn = TOPBIT; ptn != *you && *own >= bit; ptn >>= 1)
				bit <<= 1;
			if (*own > bit)
				return;
			if (*own < bit)
				break;
		}
		if (own > BOARDE) {
			//COUNT4++;

			C4[threadIdx.x] = C4[threadIdx.x] + 4;

			// printf("4\n");
			return;
		}
	}
	// 270-degree rotation
	if (*BOARD1 == TOPBIT) {
		for (ptn = TOPBIT >> 1, own = BOARD + 1; own <= BOARDE;
				own++, ptn >>= 1) {
			bit = 1;
			for (you = BOARD; *you != ptn && *own >= bit; you++)
				bit <<= 1;
			if (*own > bit)
				return;
			if (*own < bit)
				break;
		}
	}
	//COUNT8++;

	C8[threadIdx.x] = C8[threadIdx.x] + 8;

	//printf("8\n");

}

/**********************************************/
/* Primer reina interna                       */
/**********************************************/
__device__ void Backtrack2(unsigned long y, unsigned long left,
		unsigned long down, unsigned long right, int SIZE, unsigned int BOUND1,
		unsigned long *BOARD, unsigned long SIDEMASK, unsigned long *BOARDE,
		unsigned long TOPBIT, unsigned long ENDBIT, unsigned long *BOARD2,
		unsigned long *BOARD1, int SIZEE, unsigned long LASTMASK,
		unsigned long BOUND2, unsigned long MASK) {
	//printf("y %d, left %d, down %d, right %d, SIZE 	%d, BOUND1 %d, sidemask %d, topbit %d, endbit %d, size %d, lastmask %d, bound2 %d, mask %d \n", y,  left,  down,  right,  SIZE,   BOUND1,   SIDEMASK, TOPBIT,  ENDBIT,  SIZEE,  LASTMASK,  BOUND2,  MASK);
	unsigned long bitmap, bit;
	bitmap = MASK & ~(left | down | right);
	if (y == SIZEE) {
		if (bitmap) {
			if (!(bitmap & LASTMASK)) {
				//printf(" Backtrack2 => board[%d] = %lld \n",y,BOARD[y]);
				BOARD[y] = bitmap;

				Check(BOARD, BOARDE, TOPBIT, ENDBIT, BOARD2, BOARD1);
			}
		}
	} else {
		if (y < BOUND1) {
			bitmap |= SIDEMASK;
			bitmap ^= SIDEMASK;
		} else if (y == BOUND2) {
			if (!(down & SIDEMASK))
				return;
			if ((down & SIDEMASK) != SIDEMASK)
				bitmap &= SIDEMASK;
		}
		while (bitmap) {
			//  printf(" Backtrack2 => board[%d] = %lld \n",y,BOARD[y]);
			bitmap ^= BOARD[y] = bit = -bitmap & bitmap;
			Backtrack2(y + 1, (left | bit) << 1, down | bit, (right | bit) >> 1,
					SIZE, BOUND1, BOARD, SIDEMASK, BOARDE, TOPBIT, ENDBIT,
					BOARD2, BOARD1, SIZEE, LASTMASK, BOUND2, MASK);
		}
	}
}

__device__ void Backtrack2Fijo(unsigned long y, unsigned long left,
		unsigned long down, unsigned long right, unsigned long *otros, int SIZE,
		unsigned long SIDEMASK, unsigned int BOUND1, unsigned long *BOARD,
		unsigned long *BOARDE, unsigned long TOPBIT, unsigned long ENDBIT,
		unsigned long *BOARD2, unsigned long *BOARD1, int SIZEE,
		unsigned long LASTMASK, unsigned long BOUND2, int CFIJA,
		unsigned long MASK) {

	unsigned long bitmap, bit;

	bitmap = MASK & ~(left | down | right);
	if (y < BOUND1) {
		bitmap |= SIDEMASK;
		bitmap ^= SIDEMASK;
	} else if (y == BOUND2) {
		if (!(down & SIDEMASK))
			return;
		if ((down & SIDEMASK) != SIDEMASK)
			bitmap &= SIDEMASK;
	};
	//printf("Backtrack2Fijo => board[%d] = %lld \n",y,BOARD[y]);
	BOARD[y] = bit = 1 << otros[y];
	if (bitmap & bit) {
		if (y < (CFIJA - 1)) {
			Backtrack2Fijo(y + 1, (left | bit) << 1, down | bit,
					(right | bit) >> 1, otros, SIZE, SIDEMASK, BOUND1, BOARD,
					BOARDE, TOPBIT, ENDBIT, BOARD2, BOARD1, SIZEE, LASTMASK,
					BOUND2, CFIJA, MASK);
		} else
			Backtrack2(y + 1, (left | bit) << 1, down | bit, (right | bit) >> 1,
					SIZE, BOUND1, BOARD, SIDEMASK, BOARDE, TOPBIT, ENDBIT,
					BOARD2, BOARD1, SIZEE, LASTMASK, BOUND2, MASK);
	}
}
;

/**********************************************/
/* Primer reina en la esquina                 */
/**********************************************/
__device__ void Backtrack1(unsigned long y, unsigned long left,
		unsigned long down, unsigned long right, int SIZE, unsigned int BOUND1,
		unsigned long *BOARD, int SIZEE, unsigned long MASK) {

	unsigned long bitmap, bit;
	bitmap = MASK & ~(left | down | right);
	if (y == SIZEE) {
		if (bitmap) {
			// printf(" Backtrack1 => board[%d] = %lld \n",y,BOARD[y]);
			BOARD[y] = bitmap;
			//cuenta las combinaciones
			//COUNT8++;

			C8[threadIdx.x] = C8[threadIdx.x] + 8;
			//printf("8\n");

		}
	} else {
		if (y < BOUND1) {
			bitmap |= 2;
			bitmap ^= 2;
		}
		while (bitmap) {
			// printf("Backtrack1 => board[%d] = %lld \n",y,BOARD[y]);
			bitmap ^= BOARD[y] = bit = -bitmap & bitmap;
			Backtrack1(y + 1, (left | bit) << 1, down | bit, (right | bit) >> 1,
					SIZE, BOUND1, BOARD, SIZEE, MASK);
		}
	}
}

__device__ void Backtrack1Fijo(unsigned long y, unsigned long left,
		unsigned long down, unsigned long right, unsigned long *otros, int SIZE,
		unsigned long long BOUND1, unsigned long *BOARD, int CFIJA, int SIZEE,
		unsigned long MASK) {

	unsigned long bitmap, bit;

	bitmap = MASK & ~(left | down | right);
	if (y < BOUND1) {
		bitmap |= 2;
		bitmap ^= 2;
	};
	//printf(" Backtrack1Fijo => board[%d] =  %lld \n",y,BOARD[y]);
	BOARD[y] = bit = 1 << otros[y];
	if (bitmap & bit)
		if (y < (CFIJA - 1)) {
			Backtrack1Fijo(y + 1, (left | bit) << 1, down | bit,
					(right | bit) >> 1, otros, SIZE, BOUND1, BOARD, CFIJA,
					SIZEE, MASK);
		} else
			Backtrack1(y + 1, (left | bit) << 1, down | bit, (right | bit) >> 1,
					SIZE, BOUND1, BOARD, SIZEE, MASK);
}
;

/**********************************************/
/* Procesar una combinacion                   */
/**********************************************/
//- No se pueden ejecutar varios kernels en paralelo en el mismo device.
__global__ void procesoNodo(int cantCombinaciones, int CFIJA, int SIZE,
		unsigned long *resultado_gpu) {
	//VARIABLES Q ESTABAN DEFINIDAS GLOBALMENTE
	int tidx = threadIdx.x + blockIdx.x * blockDim.x;
	//
	//printf("combinacion nro %d, threadId %lld \n",combinacionAResolver, tidx);
	while (tidx < cantCombinaciones) { //los demas Threads del bloque no hacen nada, si ya llegamos al limite
		unsigned long BOUND2, BOUND1;
		unsigned long SIDEMASK;
		unsigned long ENDBIT;
		unsigned long LASTMASK;
		unsigned long *BOARD1, *BOARD2;
		unsigned long bit;
		int com;
		int SIZEE = SIZE - 1;
		unsigned long TOPBIT;

		//printf("Bloque %d \n",bidx);

		/*****/
		unsigned long MASK;
		unsigned long *BOARD;
		unsigned long *BOARDE;
		BOARD = (unsigned long *) malloc(sizeof(long) * SIZE);

		SIZEE = SIZE - 1;
		BOARDE = &BOARD[SIZEE];

		//esta variable es utilizada por los metodos dentro de la GPU creo q no es necesaria aca
		MASK = (1 << SIZE) - 1; //despalzamiento hacia la derecha
		TOPBIT = 1 << SIZEE; //desplazamiento hacia la derecha
		SIDEMASK = TOPBIT | 1; //Operación con OR con bits

		unsigned long *valores;

		//printf("thread id %lld del bloque %d y la combinacion a resolver es:%lld  \n", tidx,blockIdx.x, combinacionAResolver);
		valores = (unsigned long*) malloc(sizeof(unsigned long) * CFIJA);

		//Determina las posiciones fijas para cada fila
		//CFIJA => CANTIDAD DE FILAS PARA FORMAR LAS COMBINACIONES
		//SIZE =>  TAMAÑO DEL TABLERO
		//com => ENTERO Q INDICA EL NRO DE COMBINACION Q LE CORRESPONDE RESOLVER EN ESTA OPORTUNIDAD
		com = tidx;
		for (int i = 1; i < CFIJA; i++) {
			//% = MOD =>
			valores[i] = com % SIZE;
			com = com / SIZE;
		};
		valores[0] = com;
		SIDEMASK = TOPBIT | 1;

		if (valores[0] == 0) { //Si la Primer reina esta en la esquina
			if ((valores[1] > 1) && (valores[1] < SIZEE)) {
				BOARD[0] = 1;
				BOARD[1] = bit = 1 << valores[1];
				BOUND1 = valores[1];
				// siguiente_fila/diagonal_izquierda_ocupada/columna_ocupada/diagonal_derecha_ocupada
				Backtrack1Fijo(2, (2 | bit) << 1, 1 | bit, bit >> 1, valores,
						SIZE, BOUND1, BOARD, CFIJA, SIZEE, MASK);
			}; //La combinacioN ya es invalida!!!
		} else { //Si la Primer reina esta en el interior
			LASTMASK = SIDEMASK;
			for (int i = 1; i < valores[0]; i++) {
				LASTMASK |= LASTMASK >> 1 | LASTMASK << 1;
			}
			BOUND1 = valores[0];
			ENDBIT = TOPBIT >> BOUND1;
			BOUND2 = SIZEE - valores[0];
			BOARD1 = &BOARD[BOUND1];
			BOARD2 = &BOARD[BOUND2];
			BOARD[0] = bit = 1 << BOUND1;
			Backtrack2Fijo(1, bit << 1, bit, bit >> 1, valores, SIZE, SIDEMASK,
					BOUND1, BOARD, BOARDE, TOPBIT, ENDBIT, BOARD2, BOARD1,
					SIZEE, LASTMASK, BOUND2, CFIJA, MASK);
		}
		free(valores);
		tidx = tidx + (blockDim.x * gridDim.x);
	}

	if ((threadIdx.x + blockIdx.x * blockDim.x) < cantCombinaciones) {
		__syncthreads();

		//Agrupacion por bloque
		if (threadIdx.x == 0) {
			for (int i = 0; i < 1024; i++) {

				if (C2[i] != 0) {

					resultado_gpu[blockIdx.x] = resultado_gpu[blockIdx.x]
							+ C2[i];
					//printf("bloque %d, %lld,  %lld \n",blockIdx.x,C2[i]*2,resultado_gpu[blockIdx.x]);
				}
				if (C4[i] != 0) {
					resultado_gpu[blockIdx.x] = resultado_gpu[blockIdx.x]
							+ C4[i];
					//printf("bloque %d , %lld, %lld \n",blockIdx.x,C4[i]*4,resultado_gpu[blockIdx.x]);
				}

				if (C8[i] != 0) {
					resultado_gpu[blockIdx.x] = resultado_gpu[blockIdx.x]
							+ C8[i];
					//printf("bloque %d, %lld , %lld \n",blockIdx.x,C8[i]*8,resultado_gpu[blockIdx.x]);
				}
			}
		}
	}

}

/**************************************************************************/
/* Programa principal (MAIN)                                              */
/**************************************************************************/
int main(int argc, char *argv[]) {
	cudaEvent_t start, stop;
	float time;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	int threadPorBloque = 256;
	int CFIJA, SIZE;
	unsigned long *resultado_cpu = 0;
	unsigned long *resultado_gpu = 0;
	int i, cantComb;
	unsigned long TOTAL = 0;
	int bloques;
	int threads =256;

	// argv[1] => tamaño del tablero
	SIZE = atoi(argv[1]);
	// argv[2] => Cantidad de filas para formar las combinaciones
	CFIJA = atoi(argv[2]);

	//Determina la cantidad total de combinaciones en base a la formula:  W = (N^f)/2
	// como parámetro y que debe resolver el proceso
	cantComb = SIZE / 2;

	//Ahora eleva a la potencia de acuerdo a la cantidad de filas
	for (i = 1; i < CFIJA; i++) {
		cantComb *= SIZE;
	}
	//CantComb equivale a la cantidad de Threds en TOTAL!
	printf(
			"cantidad de Threads necesarios en total para procesar las combinaciones es: %d \n",
			cantComb);

	//Esta Tarea la realiza el thread i correspondiente al bloque i
	//            B / T
	bloques = 20;
	threads = cantComb / bloques;

	//esta variable contendra los resultados parciales de los bloques
	resultado_cpu = (unsigned long*) malloc(bloques * sizeof(long));

	//Alocamos memoria para el totalizador de soluciones
	CudaSafeCall(cudaMalloc((void**)&resultado_gpu,bloques* sizeof(long)));

	cudaMemset(resultado_gpu, 0, bloques * sizeof(long));
	/***SETEAMOS LLOS NUEVOS VALORES PARA LA STACK Y HEAP*****/

	unsigned long long limiteNew = 5000;
	// cudaDeviceSetLimit(cudaLimitMallocHeapSize,limiteNew*sizeof(long long));
	cudaDeviceSetLimit(cudaLimitStackSize, limiteNew * sizeof(long long));

	/*******llamamos al proceso********/

	printf("Para un tablero de n= %d - Bloques %d , Thread %d \n", SIZE,
			bloques, threads);
	cudaEventRecord(start, 0);
	procesoNodo<<<bloques, 512>>>(cantComb, CFIJA, SIZE, resultado_gpu);
	cudaThreadSynchronize();
	CudaCheckError();
	cudaEventRecord(stop, 0);
	cudaEventElapsedTime(&time, start, stop);

	/************************************/

	CudaSafeCall(
			cudaMemcpy(resultado_cpu,resultado_gpu,bloques* sizeof(long), cudaMemcpyDeviceToHost));
	cudaFree(resultado_gpu);

	for (int i = 0; i < bloques; i++) {
		TOTAL += resultado_cpu[i];
	}
	printf("cant soluciones %ld \n", TOTAL);

	printf("Time for the kernel: %f s \n", time / 1000);

}
;

